#ifndef _SVM_KERNEL_
#define _SVM_KERNEL_

#include "reduce.h"
#include <math.h>

//compute phi(A'A)
static __device__ __host__ float selfKernel(float *pointerA, int n_dim) {
    float acc = 0.0f;
	float value;

	for(int i=0; i< n_dim; i++){
		value = *(pointerA+i);
		acc += value * value;
	}

	return acc;
}
    
//compute phi(A'B)
static __device__ __host__ float kernel(float *pointerA, float *pointerB, int n_dim) {
    float acc = 0.0f;
	for(int i=0; i<n_dim; i++)
		acc += (*(pointerA+i)) * (*(pointerB+i));

	return acc;
}


//compute phi(a,b) and phi(a,c), where b and c are already in shared memory
//and a is in global memory
static __device__ __host__ void dualKernel(float *pointerA, float *pointerB, float *pointerC, int n_dim, float& phi1, float& phi2) {
    float acc1 = 0.0f; float acc2 = 0.0f;
    for(int i=0; i<n_dim; i++){
        float a = *(pointerA+i);
        acc1 += a * (*(pointerB+i));
        acc2 += a * (*(pointerC+i));
    } 
    phi1 = acc1; phi2 = acc2;
}


#endif
